/* MD5 code for Multiforcer */

/* MD5 Defines as per RFC reference implementation */

typedef uint32_t UINT4;


#define MD5S11 7
#define MD5S12 12
#define MD5S13 17
#define MD5S14 22
#define MD5S21 5
#define MD5S22 9
#define MD5S23 14
#define MD5S24 20
#define MD5S31 4
#define MD5S32 11
#define MD5S33 16
#define MD5S34 23
#define MD5S41 6
#define MD5S42 10
#define MD5S43 15
#define MD5S44 21

// OPTIMIZED MD5 FUNCTIONS HERE
//#define MD5F(x,y,z) (((y ^ z) & x) ^ z)
//#define MD5G(x,y,z) (((x & y) & z) ^ y)

/* F, G, H and I are basic MD5 functions.
 */
#define MD5F(x, y, z) (((x) & (y)) | ((~x) & (z)))
#define MD5G(x, y, z) (((x) & (z)) | ((y) & (~z)))
#define MD5H(x, y, z) ((x) ^ (y) ^ (z))
#define MD5I(x, y, z) ((y) ^ ((x) | (~z)))

/* ROTATE_LEFT rotates x left n bits.
 */
#define MD5ROTATE_LEFT(x, n) (((x) << (n)) | ((x) >> (32-(n))))

/* FF, GG, HH, and II transformations for rounds 1, 2, 3, and 4.
Rotation is separate from addition to prevent recomputation.
 */
#define MD5FF(a, b, c, d, x, s, ac) { \
 (a) += MD5F ((b), (c), (d)) + (x) + (UINT4)(ac); \
 (a) = MD5ROTATE_LEFT ((a), (s)); \
 (a) += (b); \
  }
#define MD5GG(a, b, c, d, x, s, ac) { \
 (a) += MD5G ((b), (c), (d)) + (x) + (UINT4)(ac); \
 (a) = MD5ROTATE_LEFT ((a), (s)); \
 (a) += (b); \
  }
#define MD5HH(a, b, c, d, x, s, ac) { \
 (a) += MD5H ((b), (c), (d)) + (x) + (UINT4)(ac); \
 (a) = MD5ROTATE_LEFT ((a), (s)); \
 (a) += (b); \
  }
#define MD5II(a, b, c, d, x, s, ac) { \
 (a) += MD5I ((b), (c), (d)) + (x) + (UINT4)(ac); \
 (a) = MD5ROTATE_LEFT ((a), (s)); \
 (a) += (b); \
  }


#define MD5_CUDA_KERNEL_CREATE(length) \
__global__ void CUDA_MD5_Search_##length (unsigned char *OutputPassword, unsigned char *success, \
			    int charsetLen, uint32_t numberOfPasswords, struct start_positions *DEVICE_Start_Positions, unsigned int count,  \
				unsigned char * DEVICE_Hashes, unsigned char *DEVICE_HashTable) { \
  const int pass_length = length; \
  uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15; \
  uint32_t a,b,c,d; \
  uint32_t thread_index = blockIdx.x*blockDim.x + threadIdx.x; \
  uint32_t *DEVICE_Hashes_32 = (uint32_t *)DEVICE_Hashes; \
  unsigned char p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15; \
  UINT4 password_count = 0; \
  __shared__ __align__(16) unsigned char sharedCharset[4096]; \
  __shared__ __align__(16) unsigned char sharedBitmap[8192]; \
  __shared__ unsigned char sharedLengths[16]; \
  copyCharsetAndBitmap(sharedCharset, sharedBitmap, sharedLengths, charsetLen, pass_length); \
  loadStartPositions(pass_length, thread_index, DEVICE_Start_Positions,  \
		   p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15); \
  while (password_count < count) { \
  initMD(pass_length, sharedCharset, \
  	p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, \
	b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15);	 \
  CUDA_MD5(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d); \
  checkHashMulti(pass_length, sharedBitmap, DEVICE_HashTable, numberOfPasswords, \
		DEVICE_Hashes_32, success, OutputPassword, \
		p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, \
		a, b, c, d, b0, b1, b2, b3, b4, b5); \
  password_count++; \
  incrementCounters##length##Multi(); \
  } \
}

#define FASTMD5_CUDA_KERNEL_CREATE(length) \
__global__ void CUDA_FASTMD5_Search_##length (unsigned char *OutputPassword, unsigned char *success,  \
			    int charsetLen, uint32_t numberOfPasswords, struct start_positions *DEVICE_Start_Positions,  \
				unsigned int count, unsigned char * DEVICE_Hashes, unsigned char *DEVICE_HashTable) { \
  const int pass_length = length; \
  uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15; \
  uint32_t a,b,c,d; \
  uint32_t thread_index = blockIdx.x*blockDim.x + threadIdx.x; \
  uint32_t *DEVICE_Hashes_32 = (uint32_t *)deviceSearchHashes; \
  unsigned char p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15; \
  UINT4 password_count = 0; \
  __shared__ __align__(16) unsigned char sharedCharset[256]; \
  copyCharsetFAST(sharedCharset, charsetLen); \
  loadStartPositions(pass_length, thread_index, DEVICE_Start_Positions,  \
		   p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15); \
  while (password_count < count) { \
  initMDFAST(pass_length, sharedCharset, \
  	p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, \
	b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15);	 \
  CUDA_MD5(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d); \
  checkHashFAST(pass_length, numberOfPasswords, \
		DEVICE_Hashes_32, success, OutputPassword, \
		p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, \
		a, b, c, d, b0); \
  password_count++; \
  incrementCounters##length (); \
  } \
}


/*
 * MD5 code: Call as:
 * CUDA_MD5(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d); 
 * Takes inputs from b0-b15 and returns a, b, c, d.
*/
__device__ inline void CUDA_MD5(UINT4 b0, UINT4 b1, UINT4 b2, UINT4 b3, UINT4 b4, UINT4 b5, UINT4 b6, UINT4 b7, 
			   UINT4 b8, UINT4 b9, UINT4 b10, UINT4 b11, UINT4 b12, UINT4 b13, UINT4 b14, UINT4 b15, 
			   UINT4 &a, UINT4 &b, UINT4 &c, UINT4 &d) {
  a = 0x67452301;
  b = 0xefcdab89;
  c = 0x98badcfe;
  d = 0x10325476;  

  MD5FF (a, b, c, d, b0, MD5S11, 0xd76aa478); /* 1 */
  MD5FF (d, a, b, c, b1, MD5S12, 0xe8c7b756); /* 2 */
  MD5FF (c, d, a, b, b2, MD5S13, 0x242070db); /* 3 */
  MD5FF (b, c, d, a, b3, MD5S14, 0xc1bdceee); /* 4 */
  MD5FF (a, b, c, d, b4, MD5S11, 0xf57c0faf); /* 5 */
  MD5FF (d, a, b, c, b5, MD5S12, 0x4787c62a); /* 6 */
  MD5FF (c, d, a, b, b6, MD5S13, 0xa8304613); /* 7 */
  MD5FF (b, c, d, a, b7, MD5S14, 0xfd469501); /* 8 */
  MD5FF (a, b, c, d, b8, MD5S11, 0x698098d8); /* 9 */
  MD5FF (d, a, b, c, b9, MD5S12, 0x8b44f7af); /* 10 */
  MD5FF (c, d, a, b, b10, MD5S13, 0xffff5bb1); /* 11 */
  MD5FF (b, c, d, a, b11, MD5S14, 0x895cd7be); /* 12 */
  MD5FF (a, b, c, d, b12, MD5S11, 0x6b901122); /* 13 */
  MD5FF (d, a, b, c, b13, MD5S12, 0xfd987193); /* 14 */
  MD5FF (c, d, a, b, b14, MD5S13, 0xa679438e); /* 15 */
  MD5FF (b, c, d, a, b15, MD5S14, 0x49b40821); /* 16 */

 /* Round 2 */
  MD5GG (a, b, c, d, b1, MD5S21, 0xf61e2562); /* 17 */
  MD5GG (d, a, b, c, b6, MD5S22, 0xc040b340); /* 18 */
  MD5GG (c, d, a, b, b11, MD5S23, 0x265e5a51); /* 19 */
  MD5GG (b, c, d, a, b0, MD5S24, 0xe9b6c7aa); /* 20 */
  MD5GG (a, b, c, d, b5, MD5S21, 0xd62f105d); /* 21 */
  MD5GG (d, a, b, c, b10, MD5S22,  0x2441453); /* 22 */
  MD5GG (c, d, a, b, b15, MD5S23, 0xd8a1e681); /* 23 */
  MD5GG (b, c, d, a, b4, MD5S24, 0xe7d3fbc8); /* 24 */
  MD5GG (a, b, c, d, b9, MD5S21, 0x21e1cde6); /* 25 */
  MD5GG (d, a, b, c, b14, MD5S22, 0xc33707d6); /* 26 */
  MD5GG (c, d, a, b, b3, MD5S23, 0xf4d50d87); /* 27 */
  MD5GG (b, c, d, a, b8, MD5S24, 0x455a14ed); /* 28 */
  MD5GG (a, b, c, d, b13, MD5S21, 0xa9e3e905); /* 29 */
  MD5GG (d, a, b, c, b2, MD5S22, 0xfcefa3f8); /* 30 */
  MD5GG (c, d, a, b, b7, MD5S23, 0x676f02d9); /* 31 */
  MD5GG (b, c, d, a, b12, MD5S24, 0x8d2a4c8a); /* 32 */

  /* Round 3 */
  MD5HH (a, b, c, d, b5, MD5S31, 0xfffa3942); /* 33 */
  MD5HH (d, a, b, c, b8, MD5S32, 0x8771f681); /* 34 */
  MD5HH (c, d, a, b, b11, MD5S33, 0x6d9d6122); /* 35 */
  MD5HH (b, c, d, a, b14, MD5S34, 0xfde5380c); /* 36 */
  MD5HH (a, b, c, d, b1, MD5S31, 0xa4beea44); /* 37 */
  MD5HH (d, a, b, c, b4, MD5S32, 0x4bdecfa9); /* 38 */
  MD5HH (c, d, a, b, b7, MD5S33, 0xf6bb4b60); /* 39 */
  MD5HH (b, c, d, a, b10, MD5S34, 0xbebfbc70); /* 40 */
  MD5HH (a, b, c, d, b13, MD5S31, 0x289b7ec6); /* 41 */
  MD5HH (d, a, b, c, b0, MD5S32, 0xeaa127fa); /* 42 */
  MD5HH (c, d, a, b, b3, MD5S33, 0xd4ef3085); /* 43 */
  MD5HH (b, c, d, a, b6, MD5S34,  0x4881d05); /* 44 */
  MD5HH (a, b, c, d, b9, MD5S31, 0xd9d4d039); /* 45 */
  MD5HH (d, a, b, c, b12, MD5S32, 0xe6db99e5); /* 46 */
  MD5HH (c, d, a, b, b15, MD5S33, 0x1fa27cf8); /* 47 */
  MD5HH (b, c, d, a, b2, MD5S34, 0xc4ac5665); /* 48 */

  /* Round 4 */
  MD5II (a, b, c, d, b0, MD5S41, 0xf4292244); /* 49 */
  MD5II (d, a, b, c, b7, MD5S42, 0x432aff97); /* 50 */
  MD5II (c, d, a, b, b14, MD5S43, 0xab9423a7); /* 51 */
  MD5II (b, c, d, a, b5, MD5S44, 0xfc93a039); /* 52 */
  MD5II (a, b, c, d, b12, MD5S41, 0x655b59c3); /* 53 */
  MD5II (d, a, b, c, b3, MD5S42, 0x8f0ccc92); /* 54 */
  MD5II (c, d, a, b, b10, MD5S43, 0xffeff47d); /* 55 */
  MD5II (b, c, d, a, b1, MD5S44, 0x85845dd1); /* 56 */
  MD5II (a, b, c, d, b8, MD5S41, 0x6fa87e4f); /* 57 */
  MD5II (d, a, b, c, b15, MD5S42, 0xfe2ce6e0); /* 58 */
  MD5II (c, d, a, b, b6, MD5S43, 0xa3014314); /* 59 */
  MD5II (b, c, d, a, b13, MD5S44, 0x4e0811a1); /* 60 */
  MD5II (a, b, c, d, b4, MD5S41, 0xf7537e82); /* 61 */
  MD5II (d, a, b, c, b11, MD5S42, 0xbd3af235); /* 62 */
  MD5II (c, d, a, b, b2, MD5S43, 0x2ad7d2bb); /* 63 */
  MD5II (b, c, d, a, b9, MD5S44, 0xeb86d391); /* 64 */
  
  // Finally, add initial values, as this is the only pass we make.
  a += 0x67452301;
  b += 0xefcdab89;
  c += 0x98badcfe;
  d += 0x10325476;  
}

/* CPU_MD5: Fast MD5 implementation
 * toHash: 64 byte input array: MUST BE ZERORED OTHER THAN DATA TO HASH
 * length: Length, in bytes, of the input string
 * hashResult: 16 byte output array 
 */

void CPU_MD5(unsigned char *toHash, unsigned int length, unsigned char *hashResult) {
  // 32-bit unsigned values for the hash
  UINT4 a,b,c,d;
  UINT4 b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15;

  // 32-bit accesses to the hash arrays  
  UINT4 *InitialArray32;
  UINT4 *OutputArray32;
  InitialArray32 = (UINT4 *)toHash;
  OutputArray32 = (UINT4 *)hashResult;


  
  b0 = (UINT4)InitialArray32[0];
  b1 = (UINT4)InitialArray32[1];
  b2 = (UINT4)InitialArray32[2];
  b3 = (UINT4)InitialArray32[3];
  b4 = (UINT4)InitialArray32[4];
  b5 = (UINT4)InitialArray32[5];
  b6 = (UINT4)InitialArray32[6];
  b7 = (UINT4)InitialArray32[7];
  b8 = (UINT4)InitialArray32[8];
  b9 = (UINT4)InitialArray32[9];
  b10 = (UINT4)InitialArray32[10];
  b11 = (UINT4)InitialArray32[11];
  b12 = (UINT4)InitialArray32[12];
  b13 = (UINT4)InitialArray32[13];
  b14 = (UINT4)InitialArray32[14];
  
  switch (length) {
    case 0:
	  b0 |= 0x00000080;
	  break;
	case 1:
	  b0 |= 0x00008000;
	  break;
	case 2:
	  b0 |= 0x00800000;
	  break;
	case 3:
	  b0 |= 0x80000000;
	  break;
    case 4:
	  b1 |= 0x00000080;
	  break;
	case 5:
	  b1 |= 0x00008000;
	  break;
	case 6:
	  b1 |= 0x00800000;
	  break;
	case 7:
	  b1 |= 0x80000000;
	  break;
    case 8:
	  b2 |= 0x00000080;
	  break;
	case 9:
	  b2 |= 0x00008000;
	  break;
	case 10:
	  b2 |= 0x00800000;
	  break;
	case 11:
	  b2 |= 0x80000000;
	  break;
    case 12:
	  b3 |= 0x00000080;
	  break;
	case 13:
	  b3 |= 0x00008000;
	  break;
	case 14:
	  b3 |= 0x00800000;
	  break;
	case 15:
	  b3 |= 0x80000000;
	  break;
    case 16:
	  b4 |= 0x00000080;
	  break;
	case 17:
	  b4 |= 0x00008000;
	  break;
	case 18:
	  b4 |= 0x00800000;
	  break;
	case 19:
	  b4 |= 0x80000000;
	  break;

  }
  b14 = length * 8;
  b15 = 0x00000000;
  
  a = 0x67452301;
  b = 0xefcdab89;
  c = 0x98badcfe;
  d = 0x10325476;  

  MD5FF (a, b, c, d, b0, MD5S11, 0xd76aa478); /* 1 */
  MD5FF (d, a, b, c, b1, MD5S12, 0xe8c7b756); /* 2 */
  MD5FF (c, d, a, b, b2, MD5S13, 0x242070db); /* 3 */
  MD5FF (b, c, d, a, b3, MD5S14, 0xc1bdceee); /* 4 */
  MD5FF (a, b, c, d, b4, MD5S11, 0xf57c0faf); /* 5 */
  MD5FF (d, a, b, c, b5, MD5S12, 0x4787c62a); /* 6 */
  MD5FF (c, d, a, b, b6, MD5S13, 0xa8304613); /* 7 */
  MD5FF (b, c, d, a, b7, MD5S14, 0xfd469501); /* 8 */
  MD5FF (a, b, c, d, b8, MD5S11, 0x698098d8); /* 9 */
  MD5FF (d, a, b, c, b9, MD5S12, 0x8b44f7af); /* 10 */
  MD5FF (c, d, a, b, b10, MD5S13, 0xffff5bb1); /* 11 */
  MD5FF (b, c, d, a, b11, MD5S14, 0x895cd7be); /* 12 */
  MD5FF (a, b, c, d, b12, MD5S11, 0x6b901122); /* 13 */
  MD5FF (d, a, b, c, b13, MD5S12, 0xfd987193); /* 14 */
  MD5FF (c, d, a, b, b14, MD5S13, 0xa679438e); /* 15 */
  MD5FF (b, c, d, a, b15, MD5S14, 0x49b40821); /* 16 */

 /* Round 2 */
  MD5GG (a, b, c, d, b1, MD5S21, 0xf61e2562); /* 17 */
  MD5GG (d, a, b, c, b6, MD5S22, 0xc040b340); /* 18 */
  MD5GG (c, d, a, b, b11, MD5S23, 0x265e5a51); /* 19 */
  MD5GG (b, c, d, a, b0, MD5S24, 0xe9b6c7aa); /* 20 */
  MD5GG (a, b, c, d, b5, MD5S21, 0xd62f105d); /* 21 */
  MD5GG (d, a, b, c, b10, MD5S22,  0x2441453); /* 22 */
  MD5GG (c, d, a, b, b15, MD5S23, 0xd8a1e681); /* 23 */
  MD5GG (b, c, d, a, b4, MD5S24, 0xe7d3fbc8); /* 24 */
  MD5GG (a, b, c, d, b9, MD5S21, 0x21e1cde6); /* 25 */
  MD5GG (d, a, b, c, b14, MD5S22, 0xc33707d6); /* 26 */
  MD5GG (c, d, a, b, b3, MD5S23, 0xf4d50d87); /* 27 */
  MD5GG (b, c, d, a, b8, MD5S24, 0x455a14ed); /* 28 */
  MD5GG (a, b, c, d, b13, MD5S21, 0xa9e3e905); /* 29 */
  MD5GG (d, a, b, c, b2, MD5S22, 0xfcefa3f8); /* 30 */
  MD5GG (c, d, a, b, b7, MD5S23, 0x676f02d9); /* 31 */
  MD5GG (b, c, d, a, b12, MD5S24, 0x8d2a4c8a); /* 32 */

  /* Round 3 */
  MD5HH (a, b, c, d, b5, MD5S31, 0xfffa3942); /* 33 */
  MD5HH (d, a, b, c, b8, MD5S32, 0x8771f681); /* 34 */
  MD5HH (c, d, a, b, b11, MD5S33, 0x6d9d6122); /* 35 */
  MD5HH (b, c, d, a, b14, MD5S34, 0xfde5380c); /* 36 */
  MD5HH (a, b, c, d, b1, MD5S31, 0xa4beea44); /* 37 */
  MD5HH (d, a, b, c, b4, MD5S32, 0x4bdecfa9); /* 38 */
  MD5HH (c, d, a, b, b7, MD5S33, 0xf6bb4b60); /* 39 */
  MD5HH (b, c, d, a, b10, MD5S34, 0xbebfbc70); /* 40 */
  MD5HH (a, b, c, d, b13, MD5S31, 0x289b7ec6); /* 41 */
  MD5HH (d, a, b, c, b0, MD5S32, 0xeaa127fa); /* 42 */
  MD5HH (c, d, a, b, b3, MD5S33, 0xd4ef3085); /* 43 */
  MD5HH (b, c, d, a, b6, MD5S34,  0x4881d05); /* 44 */
  MD5HH (a, b, c, d, b9, MD5S31, 0xd9d4d039); /* 45 */
  MD5HH (d, a, b, c, b12, MD5S32, 0xe6db99e5); /* 46 */
  MD5HH (c, d, a, b, b15, MD5S33, 0x1fa27cf8); /* 47 */
  MD5HH (b, c, d, a, b2, MD5S34, 0xc4ac5665); /* 48 */

  /* Round 4 */
  MD5II (a, b, c, d, b0, MD5S41, 0xf4292244); /* 49 */
  MD5II (d, a, b, c, b7, MD5S42, 0x432aff97); /* 50 */
  MD5II (c, d, a, b, b14, MD5S43, 0xab9423a7); /* 51 */
  MD5II (b, c, d, a, b5, MD5S44, 0xfc93a039); /* 52 */
  MD5II (a, b, c, d, b12, MD5S41, 0x655b59c3); /* 53 */
  MD5II (d, a, b, c, b3, MD5S42, 0x8f0ccc92); /* 54 */
  MD5II (c, d, a, b, b10, MD5S43, 0xffeff47d); /* 55 */
  MD5II (b, c, d, a, b1, MD5S44, 0x85845dd1); /* 56 */
  MD5II (a, b, c, d, b8, MD5S41, 0x6fa87e4f); /* 57 */
  MD5II (d, a, b, c, b15, MD5S42, 0xfe2ce6e0); /* 58 */
  MD5II (c, d, a, b, b6, MD5S43, 0xa3014314); /* 59 */
  MD5II (b, c, d, a, b13, MD5S44, 0x4e0811a1); /* 60 */
  MD5II (a, b, c, d, b4, MD5S41, 0xf7537e82); /* 61 */
  MD5II (d, a, b, c, b11, MD5S42, 0xbd3af235); /* 62 */
  MD5II (c, d, a, b, b2, MD5S43, 0x2ad7d2bb); /* 63 */
  MD5II (b, c, d, a, b9, MD5S44, 0xeb86d391); /* 64 */
  
  // Finally, add initial values, as this is the only pass we make.
  a += 0x67452301;
  b += 0xefcdab89;
  c += 0x98badcfe;
  d += 0x10325476;     
  
  OutputArray32[0] = a;
  OutputArray32[1] = b;
  OutputArray32[2] = c;
  OutputArray32[3] = d;
}

int CUDA_MD5_Get_Thread_Count(int passlength) {
  int threads;
  switch (passlength) {
	case 4:
	case 5: 
	case 6: 
	  threads = 512;
	  break;
	case 7: 
	case 8: 
	  threads = 448;
	  break;
	case 9: 
  	case 10: 
	case 11: 
	  threads = 384;
	  break;
	case 12: 
	case 13: 
	case 14: 
	  threads = 320;
	  break;
  }
  return threads;
}




void CPU_MD5_Search_0(unsigned char *SearchHash, unsigned char *Output, unsigned char *success, unsigned char *global_charset, int charsetLen, int numberOfPasswords) {

  unsigned char hash[16];
  unsigned char passwordString[64];
  int password;
  memset(passwordString, 0, 64);
  CPU_MD5(passwordString, 0, hash);
    // Password created.  Check against arrays.
    for (password = 0; password < numberOfPasswords; password++) {
      // If there is a match, deal with it appropriately.
      if (memcmp(hash, &SearchHash[HASH_LENGTH*password], HASH_LENGTH) == 0) {
        success[password] = 1;
      }
    }

}

/*

__global__ void CUDA_MD5_Search_1(unsigned char *OutputPassword, unsigned char *success, 
			    int charsetLen, uint32_t numberOfPasswords, struct start_positions *DEVICE_Start_Positions, unsigned int count, unsigned char * DEVICE_Hashes, unsigned char *DEVICE_HashTable) {
  
  const int pass_length = 1;
  uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15;
  uint32_t a,b,c,d;
  uint32_t thread_index = blockIdx.x*blockDim.x + threadIdx.x;
  uint32_t *DEVICE_Hashes_32 = (uint32_t *)DEVICE_Hashes;
  unsigned char p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15;
  UINT4 password_count = 0;
  __shared__ __align__(16) unsigned char sharedCharset[4096];
  __shared__ __align__(16) unsigned char sharedBitmap[8192];
  __shared__ unsigned char sharedLengths[16];
  copyCharsetAndBitmap(sharedCharset, sharedBitmap, sharedLengths, charsetLen, pass_length);
  loadStartPositions(pass_length, thread_index, DEVICE_Start_Positions, 
		   p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15);
  while (password_count < count) {
  initMD(pass_length, sharedCharset,
  	p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
	b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15);	
  CUDA_MD5(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d);
  checkHashMulti(pass_length, sharedBitmap, DEVICE_HashTable, numberOfPasswords,
		DEVICE_Hashes_32, success, OutputPassword,
		p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
		a, b, c, d, b0, b1, b2, b3, b4, b5);
  password_count++;
  incrementCounters1Multi();
  }
}

__global__ void CUDA_MD5_Search_2(unsigned char *OutputPassword, unsigned char *success, 
			    int charsetLen, uint32_t numberOfPasswords, struct start_positions *DEVICE_Start_Positions, unsigned int count, unsigned char * DEVICE_Hashes, unsigned char *DEVICE_HashTable) {
  
  const int pass_length = 2;
  
  // Hash work space: 16 by 32-bit registers
  uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15;
  // Output space: 4 by 32-bit registers
  uint32_t a,b,c,d;
  // Thread index, for start position loads
  uint32_t thread_index = blockIdx.x*blockDim.x + threadIdx.x;
  
  uint32_t *DEVICE_Hashes_32 = (uint32_t *)DEVICE_Hashes;
  
  // Password charset counters
  unsigned char p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15;
  UINT4 password_count = 0;
  __shared__ __align__(16) unsigned char sharedCharset[4096];
  __shared__ __align__(16) unsigned char sharedBitmap[8192];
  __shared__ unsigned char sharedLengths[16];
  
  copyCharsetAndBitmap(sharedCharset, sharedBitmap, sharedLengths, charsetLen, pass_length);

  // Load start positions from global memory into device registers
  loadStartPositions(pass_length, thread_index, DEVICE_Start_Positions, 
		   p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15);
  
  while (password_count < count) {

  //initMD(pass_length, (unsigned char *)deviceCharset,
  initMD(pass_length, sharedCharset,
  	p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
	b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15);	
#ifdef __DEVICE_EMULATION__
  printf("Thread %d: Indexes: %d%d%d%d%d%d\n",thread_index, p0, p1, p2, p3, p4, p5);
  printf("Password: %c%c%c%c%c%c\n", sharedCharset[p0],  sharedCharset[p1 + 256],  sharedCharset[p2 + 512],  sharedCharset[p3 + 768],  sharedCharset[p4 + 1024],  sharedCharset[p5 + 1280]);
  printf("%08X %08X\n", b0, b1);
#endif
  // Hash function, inlined.
  CUDA_MD5(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d);

  // Check for valid results.
  
  checkHashMulti(pass_length, sharedBitmap, DEVICE_HashTable, numberOfPasswords,
		DEVICE_Hashes_32, success, OutputPassword,
		p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
		a, b, c, d, b0, b1, b2, b3, b4, b5);
  password_count++;
  
  incrementCounters2Multi();
  }
}

__global__ void CUDA_MD5_Search_3(unsigned char *OutputPassword, unsigned char *success, 
			    int charsetLen, uint32_t numberOfPasswords, struct start_positions *DEVICE_Start_Positions, unsigned int count, unsigned char * DEVICE_Hashes, unsigned char *DEVICE_HashTable) {
  
  const int pass_length = 3;
  
  // Hash work space: 16 by 32-bit registers
  uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15;
  // Output space: 4 by 32-bit registers
  uint32_t a,b,c,d;
  // Thread index, for start position loads
  uint32_t thread_index = blockIdx.x*blockDim.x + threadIdx.x;
  
  uint32_t *DEVICE_Hashes_32 = (uint32_t *)DEVICE_Hashes;
  
  // Password charset counters
  unsigned char p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15;
  UINT4 password_count = 0;
  __shared__ __align__(16) unsigned char sharedCharset[4096];
  __shared__ __align__(16) unsigned char sharedBitmap[8192];
  __shared__ unsigned char sharedLengths[16];
  
  copyCharsetAndBitmap(sharedCharset, sharedBitmap, sharedLengths, charsetLen, pass_length);

  // Load start positions from global memory into device registers
  loadStartPositions(pass_length, thread_index, DEVICE_Start_Positions, 
		   p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15);
  
  while (password_count < count) {

  //initMD(pass_length, (unsigned char *)deviceCharset,
  initMD(pass_length, sharedCharset,
  	p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
	b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15);	
#ifdef __DEVICE_EMULATION__
  printf("Thread %d: Indexes: %d%d%d%d%d%d\n",thread_index, p0, p1, p2, p3, p4, p5);
  printf("Password: %c%c%c%c%c%c\n", sharedCharset[p0],  sharedCharset[p1 + 256],  sharedCharset[p2 + 512],  sharedCharset[p3 + 768],  sharedCharset[p4 + 1024],  sharedCharset[p5 + 1280]);
  printf("%08X %08X\n", b0, b1);
#endif
  // Hash function, inlined.
  CUDA_MD5(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d);

  // Check for valid results.
  
  checkHashMulti(pass_length, sharedBitmap, DEVICE_HashTable, numberOfPasswords,
		DEVICE_Hashes_32, success, OutputPassword,
		p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
		a, b, c, d, b0, b1, b2, b3, b4, b5);
  password_count++;
  
  incrementCounters3Multi();
  }
}

__global__ void CUDA_MD5_Search_4(unsigned char *OutputPassword, unsigned char *success, 
			    int charsetLen, uint32_t numberOfPasswords, struct start_positions *DEVICE_Start_Positions, unsigned int count, unsigned char * DEVICE_Hashes, unsigned char *DEVICE_HashTable) {
  
  const int pass_length = 4;
  
  // Hash work space: 16 by 32-bit registers
  uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15;
  // Output space: 4 by 32-bit registers
  uint32_t a,b,c,d;
  // Thread index, for start position loads
  uint32_t thread_index = blockIdx.x*blockDim.x + threadIdx.x;
  
  uint32_t *DEVICE_Hashes_32 = (uint32_t *)DEVICE_Hashes;
  
  // Password charset counters
  unsigned char p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15;
  UINT4 password_count = 0;
  __shared__ __align__(16) unsigned char sharedCharset[4096];
  __shared__ __align__(16) unsigned char sharedBitmap[8192];
  __shared__ unsigned char sharedLengths[16];
  
  copyCharsetAndBitmap(sharedCharset, sharedBitmap, sharedLengths, charsetLen, pass_length);

  // Load start positions from global memory into device registers
  loadStartPositions(pass_length, thread_index, DEVICE_Start_Positions, 
		   p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15);
  
  while (password_count < count) {

  //initMD(pass_length, (unsigned char *)deviceCharset,
  initMD(pass_length, sharedCharset,
  	p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
	b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15);	
#ifdef __DEVICE_EMULATION__
  printf("Thread %d: Indexes: %d%d%d%d%d%d\n",thread_index, p0, p1, p2, p3, p4, p5);
  printf("Password: %c%c%c%c%c%c\n", sharedCharset[p0],  sharedCharset[p1 + 256],  sharedCharset[p2 + 512],  sharedCharset[p3 + 768],  sharedCharset[p4 + 1024],  sharedCharset[p5 + 1280]);
  printf("%08X %08X\n", b0, b1);
#endif
  // Hash function, inlined.
  CUDA_MD5(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d);

  // Check for valid results.
  
  checkHashMulti(pass_length, sharedBitmap, DEVICE_HashTable, numberOfPasswords,
		DEVICE_Hashes_32, success, OutputPassword,
		p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
		a, b, c, d, b0, b1, b2, b3, b4, b5);
  password_count++;
  
  incrementCounters4Multi();
  }
}



__global__ void CUDA_MD5_Search_5(unsigned char *OutputPassword, unsigned char *success, 
			    int charsetLen, uint32_t numberOfPasswords, struct start_positions *DEVICE_Start_Positions, unsigned int count, unsigned char * DEVICE_Hashes, unsigned char *DEVICE_HashTable) {
  
  const int pass_length = 5;
  
  // Hash work space: 16 by 32-bit registers
  uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15;
  // Output space: 4 by 32-bit registers
  uint32_t a,b,c,d;
  // Thread index, for start position loads
  uint32_t thread_index = blockIdx.x*blockDim.x + threadIdx.x;
  
  uint32_t *DEVICE_Hashes_32 = (uint32_t *)DEVICE_Hashes;
  
  // Password charset counters
  unsigned char p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15;
  UINT4 password_count = 0;
  __shared__ __align__(16) unsigned char sharedCharset[4096];
  __shared__ __align__(16) unsigned char sharedBitmap[8192];
  __shared__ unsigned char sharedLengths[16];
  
  copyCharsetAndBitmap(sharedCharset, sharedBitmap, sharedLengths, charsetLen, pass_length);

  // Load start positions from global memory into device registers
  loadStartPositions(pass_length, thread_index, DEVICE_Start_Positions, 
		   p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15);
  
  while (password_count < count) {

  //initMD(pass_length, (unsigned char *)deviceCharset,
  initMD(pass_length, sharedCharset,
  	p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
	b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15);	
#ifdef __DEVICE_EMULATION__
  printf("Thread %d: Indexes: %d%d%d%d%d%d\n",thread_index, p0, p1, p2, p3, p4, p5);
  printf("Password: %c%c%c%c%c%c\n", sharedCharset[p0],  sharedCharset[p1 + 256],  sharedCharset[p2 + 512],  sharedCharset[p3 + 768],  sharedCharset[p4 + 1024],  sharedCharset[p5 + 1280]);
  printf("%08X %08X\n", b0, b1);
#endif
  // Hash function, inlined.
  CUDA_MD5(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d);

  // Check for valid results.
  
  checkHashMulti(pass_length, sharedBitmap, DEVICE_HashTable, numberOfPasswords,
		DEVICE_Hashes_32, success, OutputPassword,
		p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
		a, b, c, d, b0, b1, b2, b3, b4, b5);
  password_count++;
  
  incrementCounters5Multi();
  }
}


__global__ void CUDA_MD5_Search_6(unsigned char *OutputPassword, unsigned char *success, 
			    int charsetLen, uint32_t numberOfPasswords, struct start_positions *DEVICE_Start_Positions, unsigned int count, unsigned char * DEVICE_Hashes, unsigned char *DEVICE_HashTable) {
  
  const int pass_length = 6;
  
  // Hash work space: 16 by 32-bit registers
  uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15;
  // Output space: 4 by 32-bit registers
  uint32_t a,b,c,d;
  // Thread index, for start position loads
  uint32_t thread_index = blockIdx.x*blockDim.x + threadIdx.x;
  
  uint32_t *DEVICE_Hashes_32 = (uint32_t *)DEVICE_Hashes;
  
  // Password charset counters
  unsigned char p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15;
  UINT4 password_count = 0;
  __shared__ __align__(16) unsigned char sharedCharset[4096];
  __shared__ __align__(16) unsigned char sharedBitmap[8192];
  __shared__ unsigned char sharedLengths[16];
  
  copyCharsetAndBitmap(sharedCharset, sharedBitmap, sharedLengths, charsetLen, pass_length);

  // Load start positions from global memory into device registers
  loadStartPositions(pass_length, thread_index, DEVICE_Start_Positions, 
		   p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15);
  
  while (password_count < count) {

  //initMD(pass_length, (unsigned char *)deviceCharset,
  initMD(pass_length, sharedCharset,
  	p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
	b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15);	
#ifdef __DEVICE_EMULATION__
  printf("Thread %d: Indexes: %d%d%d%d%d%d\n",thread_index, p0, p1, p2, p3, p4, p5);
  printf("Password: %c%c%c%c%c%c\n", sharedCharset[p0],  sharedCharset[p1 + 256],  sharedCharset[p2 + 512],  sharedCharset[p3 + 768],  sharedCharset[p4 + 1024],  sharedCharset[p5 + 1280]);
  printf("%08X %08X\n", b0, b1);
#endif
  // Hash function, inlined.
  CUDA_MD5(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d);

  // Check for valid results.
  
  checkHashMulti(pass_length, sharedBitmap, DEVICE_HashTable, numberOfPasswords,
		DEVICE_Hashes_32, success, OutputPassword,
		p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
		a, b, c, d, b0, b1, b2, b3, b4, b5);
  password_count++;
  
  incrementCounters6Multi();
  }
}


__global__ void CUDA_MD5_Search_7(unsigned char *OutputPassword, unsigned char *success, 
			    int charsetLen, uint32_t numberOfPasswords, struct start_positions *DEVICE_Start_Positions, unsigned int count, unsigned char * DEVICE_Hashes, unsigned char *DEVICE_HashTable) {
  
  
  const int pass_length = 7;
  
  // Hash work space: 16 by 32-bit registers
  uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15;
  // Output space: 4 by 32-bit registers
  uint32_t a,b,c,d;
  // Thread index, for start position loads
  uint32_t thread_index = blockIdx.x*blockDim.x + threadIdx.x;
  
  uint32_t *DEVICE_Hashes_32 = (uint32_t *)DEVICE_Hashes;
  
  // Password charset counters
  unsigned char p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15;
  UINT4 password_count = 0;
  __shared__ __align__(16) unsigned char sharedCharset[4096];
  __shared__ __align__(16) unsigned char sharedBitmap[8192];
  __shared__ __align__(8) unsigned char sharedLengths[16];
  
  copyCharsetAndBitmap(sharedCharset, sharedBitmap, sharedLengths, charsetLen, pass_length);

  // Load start positions from global memory into device registers
  loadStartPositions(pass_length, thread_index, DEVICE_Start_Positions, 
		   p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15);
  
  while (password_count < count) {

  //initMD(pass_length, (unsigned char *)deviceCharset,
  initMD(pass_length, sharedCharset,
  	p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
	b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15);	

  // Hash function, inlined.
  CUDA_MD5(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d);

  // Check for valid results.
  
  checkHashMulti(pass_length, sharedBitmap, DEVICE_HashTable, numberOfPasswords,
		DEVICE_Hashes_32, success, OutputPassword,
		p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
		a, b, c, d, b0, b1, b2, b3, b4, b5);
  password_count++;
  
  incrementCounters7Multi();
  }
}



__global__ void CUDA_MD5_Search_8(unsigned char *OutputPassword, unsigned char *success, 
			    int charsetLen, uint32_t numberOfPasswords, struct start_positions *DEVICE_Start_Positions, unsigned int count, unsigned char * DEVICE_Hashes, unsigned char *DEVICE_HashTable) {
  
  
  const int pass_length = 8;
  
  // Hash work space: 16 by 32-bit registers
  uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15;
  // Output space: 4 by 32-bit registers
  uint32_t a,b,c,d;
  // Thread index, for start position loads
  uint32_t thread_index = blockIdx.x*blockDim.x + threadIdx.x;
  
  uint32_t *DEVICE_Hashes_32 = (uint32_t *)DEVICE_Hashes;
  
  // Password charset counters
  unsigned char p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15;
  UINT4 password_count = 0;
  __shared__ __align__(16) unsigned char sharedCharset[4096];
  __shared__ __align__(16) unsigned char sharedBitmap[8192];
  __shared__ __align__(8) unsigned char sharedLengths[16];
  
  copyCharsetAndBitmap(sharedCharset, sharedBitmap, sharedLengths, charsetLen, pass_length);

  // Load start positions from global memory into device registers
  loadStartPositions(pass_length, thread_index, DEVICE_Start_Positions, 
		   p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15);
  
  while (password_count < count) {

  //initMD(pass_length, (unsigned char *)deviceCharset,
  initMD(pass_length, sharedCharset,
  	p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
	b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15);	

  // Hash function, inlined.
  CUDA_MD5(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d);

  // Check for valid results.
  
  checkHashMulti(pass_length, sharedBitmap, DEVICE_HashTable, numberOfPasswords,
		DEVICE_Hashes_32, success, OutputPassword,
		p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
		a, b, c, d, b0, b1, b2, b3, b4, b5);
  password_count++;
  
  incrementCounters8Multi();
  }
}


__global__ void CUDA_MD5_Search_9(unsigned char *OutputPassword, unsigned char *success, 
			    int charsetLen, uint32_t numberOfPasswords, struct start_positions *DEVICE_Start_Positions, unsigned int count, unsigned char * DEVICE_Hashes, unsigned char *DEVICE_HashTable) {
  const int pass_length = 9;
  uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15;
  uint32_t a,b,c,d;
  uint32_t thread_index = blockIdx.x*blockDim.x + threadIdx.x;
  uint32_t *DEVICE_Hashes_32 = (uint32_t *)DEVICE_Hashes;
  unsigned char p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15;
  UINT4 password_count = 0;
  __shared__ __align__(16) unsigned char sharedCharset[4096];
  __shared__ __align__(16) unsigned char sharedBitmap[8192];
  __shared__ __align__(8) unsigned char sharedLengths[16];
  copyCharsetAndBitmap(sharedCharset, sharedBitmap, sharedLengths, charsetLen, pass_length);
  loadStartPositions(pass_length, thread_index, DEVICE_Start_Positions, 
		   p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15);
  while (password_count < count) {
  initMD(pass_length, sharedCharset,
  	p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
	b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15);	
  CUDA_MD5(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d);
  checkHashMulti(pass_length, sharedBitmap, DEVICE_HashTable, numberOfPasswords,
		DEVICE_Hashes_32, success, OutputPassword,
		p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
		a, b, c, d, b0, b1, b2, b3, b4, b5);
  password_count++;
  incrementCounters9Multi();
  }
}

__global__ void CUDA_MD5_Search_10(unsigned char *OutputPassword, unsigned char *success, 
			    int charsetLen, uint32_t numberOfPasswords, struct start_positions *DEVICE_Start_Positions, unsigned int count, unsigned char * DEVICE_Hashes, unsigned char *DEVICE_HashTable) {
  const int pass_length = 10;
  uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15;
  uint32_t a,b,c,d;
  uint32_t thread_index = blockIdx.x*blockDim.x + threadIdx.x;
  uint32_t *DEVICE_Hashes_32 = (uint32_t *)DEVICE_Hashes;
  unsigned char p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15;
  UINT4 password_count = 0;
  __shared__ __align__(16) unsigned char sharedCharset[4096];
  __shared__ __align__(16) unsigned char sharedBitmap[8192];
  __shared__ __align__(8) unsigned char sharedLengths[16];
  copyCharsetAndBitmap(sharedCharset, sharedBitmap, sharedLengths, charsetLen, pass_length);
  loadStartPositions(pass_length, thread_index, DEVICE_Start_Positions, 
		   p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15);
  while (password_count < count) {
  initMD(pass_length, sharedCharset,
  	p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
	b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15);	
  CUDA_MD5(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d);
  checkHashMulti(pass_length, sharedBitmap, DEVICE_HashTable, numberOfPasswords,
		DEVICE_Hashes_32, success, OutputPassword,
		p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
		a, b, c, d, b0, b1, b2, b3, b4, b5);
  password_count++;
  incrementCounters10Multi();
  }
}

__global__ void CUDA_MD5_Search_11(unsigned char *OutputPassword, unsigned char *success, 
			    int charsetLen, uint32_t numberOfPasswords, struct start_positions *DEVICE_Start_Positions, unsigned int count, unsigned char * DEVICE_Hashes, unsigned char *DEVICE_HashTable) {
  const int pass_length = 11;
  uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15;
  uint32_t a,b,c,d;
  uint32_t thread_index = blockIdx.x*blockDim.x + threadIdx.x;
  uint32_t *DEVICE_Hashes_32 = (uint32_t *)DEVICE_Hashes;
  unsigned char p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15;
  UINT4 password_count = 0;
  __shared__ __align__(16) unsigned char sharedCharset[4096];
  __shared__ __align__(16) unsigned char sharedBitmap[8192];
  __shared__ __align__(8) unsigned char sharedLengths[16];
  copyCharsetAndBitmap(sharedCharset, sharedBitmap, sharedLengths, charsetLen, pass_length);
  loadStartPositions(pass_length, thread_index, DEVICE_Start_Positions, 
		   p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15);
  while (password_count < count) {
  initMD(pass_length, sharedCharset,
  	p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
	b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15);	
  CUDA_MD5(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d);
  checkHashMulti(pass_length, sharedBitmap, DEVICE_HashTable, numberOfPasswords,
		DEVICE_Hashes_32, success, OutputPassword,
		p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
		a, b, c, d, b0, b1, b2, b3, b4, b5);
  password_count++;
  incrementCounters11Multi();
  }
}

__global__ void CUDA_MD5_Search_12(unsigned char *OutputPassword, unsigned char *success, 
			    int charsetLen, uint32_t numberOfPasswords, struct start_positions *DEVICE_Start_Positions, unsigned int count, unsigned char * DEVICE_Hashes, unsigned char *DEVICE_HashTable) {
  const int pass_length = 12;
  uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15;
  uint32_t a,b,c,d;
  uint32_t thread_index = blockIdx.x*blockDim.x + threadIdx.x;
  uint32_t *DEVICE_Hashes_32 = (uint32_t *)DEVICE_Hashes;
  unsigned char p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15;
  UINT4 password_count = 0;
  __shared__ __align__(16) unsigned char sharedCharset[4096];
  __shared__ __align__(16) unsigned char sharedBitmap[8192];
  __shared__ __align__(8) unsigned char sharedLengths[16];
  copyCharsetAndBitmap(sharedCharset, sharedBitmap, sharedLengths, charsetLen, pass_length);
  loadStartPositions(pass_length, thread_index, DEVICE_Start_Positions, 
		   p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15);
  while (password_count < count) {
  initMD(pass_length, sharedCharset,
  	p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
	b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15);	
  CUDA_MD5(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d);
  checkHashMulti(pass_length, sharedBitmap, DEVICE_HashTable, numberOfPasswords,
		DEVICE_Hashes_32, success, OutputPassword,
		p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
		a, b, c, d, b0, b1, b2, b3, b4, b5);
  password_count++;
  incrementCounters12Multi();
  }
}
__global__ void CUDA_MD5_Search_13(unsigned char *OutputPassword, unsigned char *success, 
			    int charsetLen, uint32_t numberOfPasswords, struct start_positions *DEVICE_Start_Positions, unsigned int count, unsigned char * DEVICE_Hashes, unsigned char *DEVICE_HashTable) {
  const int pass_length = 13;
  uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15;
  uint32_t a,b,c,d;
  uint32_t thread_index = blockIdx.x*blockDim.x + threadIdx.x;
  uint32_t *DEVICE_Hashes_32 = (uint32_t *)DEVICE_Hashes;
  unsigned char p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15;
  UINT4 password_count = 0;
  __shared__ __align__(16) unsigned char sharedCharset[4096];
  __shared__ __align__(16) unsigned char sharedBitmap[8192];
  __shared__ __align__(8) unsigned char sharedLengths[16];
  copyCharsetAndBitmap(sharedCharset, sharedBitmap, sharedLengths, charsetLen, pass_length);
  loadStartPositions(pass_length, thread_index, DEVICE_Start_Positions, 
		   p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15);
  while (password_count < count) {
  initMD(pass_length, sharedCharset,
  	p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
	b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15);	
  CUDA_MD5(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d);
  checkHashMulti(pass_length, sharedBitmap, DEVICE_HashTable, numberOfPasswords,
		DEVICE_Hashes_32, success, OutputPassword,
		p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
		a, b, c, d, b0, b1, b2, b3, b4, b5);
  password_count++;
  incrementCounters13Multi();
  }
}
__global__ void CUDA_MD5_Search_14(unsigned char *OutputPassword, unsigned char *success, 
			    int charsetLen, uint32_t numberOfPasswords, struct start_positions *DEVICE_Start_Positions, unsigned int count, unsigned char * DEVICE_Hashes, unsigned char *DEVICE_HashTable) {
  const int pass_length = 14;
  uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15;
  uint32_t a,b,c,d;
  uint32_t thread_index = blockIdx.x*blockDim.x + threadIdx.x;
  uint32_t *DEVICE_Hashes_32 = (uint32_t *)DEVICE_Hashes;
  unsigned char p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15;
  UINT4 password_count = 0;
  __shared__ __align__(16) unsigned char sharedCharset[4096];
  __shared__ __align__(16) unsigned char sharedBitmap[8192];
  __shared__ __align__(8) unsigned char sharedLengths[16];
  copyCharsetAndBitmap(sharedCharset, sharedBitmap, sharedLengths, charsetLen, pass_length);
  loadStartPositions(pass_length, thread_index, DEVICE_Start_Positions, 
		   p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15);
  while (password_count < count) {
  initMD(pass_length, sharedCharset,
  	p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
	b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15);	
  CUDA_MD5(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d);
  checkHashMulti(pass_length, sharedBitmap, DEVICE_HashTable, numberOfPasswords,
		DEVICE_Hashes_32, success, OutputPassword,
		p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
		a, b, c, d, b0, b1, b2, b3, b4, b5);
  password_count++;
  incrementCounters14Multi();
  }
}
__global__ void CUDA_MD5_Search_15(unsigned char *OutputPassword, unsigned char *success, 
			    int charsetLen, uint32_t numberOfPasswords, struct start_positions *DEVICE_Start_Positions, unsigned int count, unsigned char * DEVICE_Hashes, unsigned char *DEVICE_HashTable) {
  const int pass_length = 15;
  uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15;
  uint32_t a,b,c,d;
  uint32_t thread_index = blockIdx.x*blockDim.x + threadIdx.x;
  uint32_t *DEVICE_Hashes_32 = (uint32_t *)DEVICE_Hashes;
  unsigned char p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15;
  UINT4 password_count = 0;
  __shared__ __align__(16) unsigned char sharedCharset[4096];
  __shared__ __align__(16) unsigned char sharedBitmap[8192];
  __shared__ __align__(8) unsigned char sharedLengths[16];
  copyCharsetAndBitmap(sharedCharset, sharedBitmap, sharedLengths, charsetLen, pass_length);
  loadStartPositions(pass_length, thread_index, DEVICE_Start_Positions, 
		   p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15);
  while (password_count < count) {
  initMD(pass_length, sharedCharset,
  	p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
	b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15);	
  CUDA_MD5(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d);
  checkHashMulti(pass_length, sharedBitmap, DEVICE_HashTable, numberOfPasswords,
		DEVICE_Hashes_32, success, OutputPassword,
		p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
		a, b, c, d, b0, b1, b2, b3, b4, b5);
  password_count++;
  incrementCounters15Multi();
  }
}
__global__ void CUDA_MD5_Search_16(unsigned char *OutputPassword, unsigned char *success, 
			    int charsetLen, uint32_t numberOfPasswords, struct start_positions *DEVICE_Start_Positions, unsigned int count, unsigned char * DEVICE_Hashes, unsigned char *DEVICE_HashTable) {
  const int pass_length = 16;
  uint32_t b0,b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15;
  uint32_t a,b,c,d;
  uint32_t thread_index = blockIdx.x*blockDim.x + threadIdx.x;
  uint32_t *DEVICE_Hashes_32 = (uint32_t *)DEVICE_Hashes;
  unsigned char p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15;
  UINT4 password_count = 0;
  __shared__ __align__(16) unsigned char sharedCharset[4096];
  __shared__ __align__(16) unsigned char sharedBitmap[8192];
  __shared__ __align__(8) unsigned char sharedLengths[16];
  copyCharsetAndBitmap(sharedCharset, sharedBitmap, sharedLengths, charsetLen, pass_length);
  loadStartPositions(pass_length, thread_index, DEVICE_Start_Positions, 
		   p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15);
  while (password_count < count) {
  initMD(pass_length, sharedCharset,
  	p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
	b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15);	
  CUDA_MD5(b0, b1, b2, b3, b4, b5, b6, b7, b8, b9, b10, b11, b12, b13, b14, b15, a, b, c, d);
  checkHashMulti(pass_length, sharedBitmap, DEVICE_HashTable, numberOfPasswords,
		DEVICE_Hashes_32, success, OutputPassword,
		p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,
		a, b, c, d, b0, b1, b2, b3, b4, b5);
  password_count++;
  incrementCounters16Multi();
  }
}
*/



MD5_CUDA_KERNEL_CREATE(1)
MD5_CUDA_KERNEL_CREATE(2)
MD5_CUDA_KERNEL_CREATE(3)
MD5_CUDA_KERNEL_CREATE(4)
MD5_CUDA_KERNEL_CREATE(5)
MD5_CUDA_KERNEL_CREATE(6)
MD5_CUDA_KERNEL_CREATE(7)
MD5_CUDA_KERNEL_CREATE(8)
MD5_CUDA_KERNEL_CREATE(9)
MD5_CUDA_KERNEL_CREATE(10)
MD5_CUDA_KERNEL_CREATE(11)
MD5_CUDA_KERNEL_CREATE(12)
MD5_CUDA_KERNEL_CREATE(13)
MD5_CUDA_KERNEL_CREATE(14)
MD5_CUDA_KERNEL_CREATE(15)
MD5_CUDA_KERNEL_CREATE(16)

FASTMD5_CUDA_KERNEL_CREATE(1)
FASTMD5_CUDA_KERNEL_CREATE(2)
FASTMD5_CUDA_KERNEL_CREATE(3)
FASTMD5_CUDA_KERNEL_CREATE(4)
FASTMD5_CUDA_KERNEL_CREATE(5)
FASTMD5_CUDA_KERNEL_CREATE(6)
FASTMD5_CUDA_KERNEL_CREATE(7)
FASTMD5_CUDA_KERNEL_CREATE(8)
FASTMD5_CUDA_KERNEL_CREATE(9)
FASTMD5_CUDA_KERNEL_CREATE(10)
FASTMD5_CUDA_KERNEL_CREATE(11)
FASTMD5_CUDA_KERNEL_CREATE(12)
FASTMD5_CUDA_KERNEL_CREATE(13)
FASTMD5_CUDA_KERNEL_CREATE(14)
FASTMD5_CUDA_KERNEL_CREATE(15)
FASTMD5_CUDA_KERNEL_CREATE(16)


void Launch_CPU_MD5_Kernel(int passlength, unsigned char *SearchHash, unsigned char *Output, unsigned char *success, 
							unsigned char *global_charset, int charsetLen, int numberOfPasswords) {
  if (passlength == 0) {
    CPU_MD5_Search_0(SearchHash, Output, success, global_charset, charsetLen, numberOfPasswords);
  }
}


void Launch_CUDA_MD5_Kernel(int passlength, uint64_t charsetLength, int numberOfPasswords, unsigned char *DEVICE_Passwords, 
						unsigned char *DEVICE_Success, struct start_positions *DEVICE_Start_Positions, uint64_t per_step, uint64_t threads, uint64_t blocks, unsigned char *DEVICE_Hashes, unsigned char *DEVICE_Bitmap) {
    if (passlength == 1) {
	  CUDA_MD5_Search_1 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap); 
	} else if (passlength == 2) {
	  CUDA_MD5_Search_2 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);
	} else if (passlength == 3) {
	  CUDA_MD5_Search_3 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);
	} else if (passlength == 4) {
	  CUDA_MD5_Search_4 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);
	} else if (passlength == 5) {
	  CUDA_MD5_Search_5 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap); 
	} else if (passlength == 6) {
	  CUDA_MD5_Search_6 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap); 
	} else if (passlength == 7) {
	  CUDA_MD5_Search_7 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap); 
	} else if (passlength == 8) {
	  CUDA_MD5_Search_8 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);  
	} else if (passlength == 9) {
	  CUDA_MD5_Search_9 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);
	} else if (passlength == 10) {
	  CUDA_MD5_Search_10 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);
	} else if (passlength == 11) {
	  CUDA_MD5_Search_11 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);
	} else if (passlength == 12) {
	  CUDA_MD5_Search_12 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap); 
	} else if (passlength == 13) {
	  CUDA_MD5_Search_13 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap); 
	} else if (passlength == 14) {
	  CUDA_MD5_Search_14 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);
	} else if (passlength == 15) {
	  CUDA_MD5_Search_15 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);
	} else if (passlength == 16) {
	  CUDA_MD5_Search_16 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);
	}
	
	cudaError_t err = cudaGetLastError();
    if( cudaSuccess != err) 
      {
        fprintf(stderr, "Cuda error: %s.\n", cudaGetErrorString( err) );
        exit(EXIT_FAILURE);
      } 
}

void Launch_CUDA_FASTMD5_Kernel(int passlength, uint64_t charsetLength, int numberOfPasswords, unsigned char *DEVICE_Passwords, 
						unsigned char *DEVICE_Success, struct start_positions *DEVICE_Start_Positions, uint64_t per_step, uint64_t threads, uint64_t blocks, unsigned char *DEVICE_Hashes, unsigned char *DEVICE_Bitmap) {
    if (passlength == 1) {
	  CUDA_FASTMD5_Search_1 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap); 
	} else if (passlength == 2) {
	  CUDA_FASTMD5_Search_2 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);
	} else if (passlength == 3) {
	  CUDA_FASTMD5_Search_3 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);
	} else if (passlength == 4) {
	  CUDA_FASTMD5_Search_4 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);
	} else if (passlength == 5) {
	  CUDA_FASTMD5_Search_5 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap); 
	} else if (passlength == 6) {
	  CUDA_FASTMD5_Search_6 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap); 
	} else if (passlength == 7) {
	  CUDA_FASTMD5_Search_7 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap); 
	} else if (passlength == 8) {
	  CUDA_FASTMD5_Search_8 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);  
	} else if (passlength == 9) {
	  CUDA_FASTMD5_Search_9 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);
	} else if (passlength == 10) {
	  CUDA_FASTMD5_Search_10 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);
	} else if (passlength == 11) {
	  CUDA_FASTMD5_Search_11 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);
	} else if (passlength == 12) {
	  CUDA_FASTMD5_Search_12 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap); 
	} else if (passlength == 13) {
	  CUDA_FASTMD5_Search_13 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap); 
	} else if (passlength == 14) {
	  CUDA_FASTMD5_Search_14 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);
	} else if (passlength == 15) {
	  CUDA_FASTMD5_Search_15 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);
	} else if (passlength == 16) {
	  CUDA_FASTMD5_Search_16 <<< blocks, threads >>> (DEVICE_Passwords, DEVICE_Success, charsetLength, numberOfPasswords, DEVICE_Start_Positions, per_step, DEVICE_Hashes, DEVICE_Bitmap);
	}
	
	cudaError_t err = cudaGetLastError();
    if( cudaSuccess != err) 
      {
        fprintf(stderr, "Cuda error: %s.\n", cudaGetErrorString( err) );
        exit(EXIT_FAILURE);
      } 
}


